ITS: new CPU + GPU seeding vertexer by cima22 · Pull Request #15733 · AliceO2Group/AliceO2

cima22 · 2026-08-27T11:12:06Z

Adds an optional seeding vertexer that runs as a prepended tracker pass (diamond trackleting -> cells -> lines -> parallel seeding), on both the CPU and GPU traits.
Additionally, unsorted clusters are loaded to GPU and then sorted on GPU.

To set it:
ITSCATrackerParam.seedingVertexIteration=1;
ITSVertexerParam.useParallelSeeding=1;

Tuned parameters for Pb--Pb:

ITSCATrackerParam.seedingVertexIteration=1;ITSVertexerParam.useParallelSeeding=1;ITSVertexerParam.useTruthSeeding=0;ITSCATrackerParam.diamondTrackletingPVres=1.5102214480735774;ITSCATrackerParam.diamondTrackletingNSigmaCut=2.976946815605669;ITSCATrackerParam.diamondTrackletingCellDeltaTanLambdaSigma=0.005566295924057422;ITSCATrackerParam.diamondCellTanLambdaNSigma=2.5;ITSCATrackerParam.diamondTrackletingCellDeltaPhiMinPt=0.12;ITSCATrackerParam.cellLineSharedClusterCut=1;ITSVertexerParam.clusterCut=0.07;ITSVertexerParam.pairCut=0.025684202919354766;ITSVertexerParam.phiCut=0.008875191490279988;ITSVertexerParam.clusterContributorsCut=3;ITSVertexerParam.lineMinPt=0.10;ITSVertexerParam.nSigmaCut=0.0016;ITSVertexerParam.goodLineChi2Cut=9.814679066181695;ITSVertexerParam.goodLinePtCut=0.12;ITSVertexerParam.goodContributorsSignificance=0.070;ITSVertexerParam.suppressLowMultDebris=13;ITSVertexerParam.fineZWindow=0.010;ITSVertexerParam.fineMinDensity=8;ITSVertexerParam.fineMaxDrift=0.005;ITSVertexerParam.duplicateZScale=0.7;ITSVertexerParam.duplicateZCut=0.1388966993405415;

Tuned parameters for pp:

ITSCATrackerParam.seedingVertexIteration=1;ITSVertexerParam.useParallelSeeding=1;ITSVertexerParam.useTruthSeeding=0;ITSCATrackerParam.diamondTrackletingPVres=3.74344532796542;ITSCATrackerParam.diamondTrackletingNSigmaCut=5.352263543811316;ITSCATrackerParam.diamondTrackletingCellDeltaTanLambdaSigma=0.002563403956702525;ITSCATrackerParam.diamondCellTanLambdaNSigma=3.257473162361808;ITSCATrackerParam.diamondTrackletingCellDeltaPhiMinPt=0.09407482537938365;ITSCATrackerParam.cellLineSharedClusterCut=3;ITSVertexerParam.clusterCut=0.05557266167329535;ITSVertexerParam.pairCut=0.05176031785438137;ITSVertexerParam.phiCut=0.017639175978851104;ITSVertexerParam.clusterContributorsCut=2;ITSVertexerParam.lineMinPt=0.1;ITSVertexerParam.nSigmaCut=0.0016;ITSVertexerParam.goodLineChi2Cut=9.814679066181695;ITSVertexerParam.goodLinePtCut=0.12;ITSVertexerParam.goodContributorsSignificance=0.07;ITSVertexerParam.suppressLowMultDebris=13;ITSVertexerParam.fineZWindow=0.013014502308220061;ITSVertexerParam.fineMinDensity=4;ITSVertexerParam.fineMaxDrift=0.004003953018135234;ITSVertexerParam.duplicateZScale=0.1230148343342053;ITSVertexerParam.duplicateZCut=0.06571512005272727;

Adds a seeding vertexer that runs as a prepended tracker pass (diamond trackleting -> cells -> lines -> parallel seeding), on both the CPU and GPU traits, replacing the per-ROF CPU vertexer for the seeding step.

shahor02 · 2026-08-27T12:14:39Z

Hi @cima22 , are there some slides or plots on the new algo., particularly its performance? Perhaps you could present it next Wed., either at the ITS meeting at 11 (if it is planned, @fprino ?) or at sim/rec WP12/13 at 15:30?

f3sch

Hi @cima22, thanks! Some preliminary comments although I have to admit to not understand the logic in its entirety yet.

f3sch · 2026-08-27T12:24:38Z

+// MC only
+template <int NLayers>
+void TimeFrameGPU<NLayers>::downloadMemberOffsetsDevice()
+{
+  const int nPeaks = mPeakOffsetsHost.empty() ? 0 : mPeakOffsetsHost.back();
+  mMemberOffsetsHost.resize(nPeaks + 1);
+  if (nPeaks) {
+    GPUChkErrS(cudaMemcpyAsync(mMemberOffsetsHost.data(), mMemberOffsetsDevice, (nPeaks + 1) * sizeof(int), cudaMemcpyDeviceToHost, mGpuStreams[0].get()));
+    GPUChkErrS(cudaStreamSynchronize(mGpuStreams[0].get()));
+  }
+  mNMembers = (nPeaks && !mMemberOffsetsHost.empty()) ? mMemberOffsetsHost.back() : 0;
+}
+
+template <int NLayers>
+void TimeFrameGPU<NLayers>::downloadMemberLinesDevice()
+{
+  mMemberLinesHost.resize(mNMembers);
+  if (mNMembers) {
+    GPUChkErrS(cudaMemcpyAsync(mMemberLinesHost.data(), mMemberLinesDevice, mNMembers * sizeof(int), cudaMemcpyDeviceToHost, mGpuStreams[0].get()));
+    GPUChkErrS(cudaStreamSynchronize(mGpuStreams[0].get()));
+  }
+}


In general, I would not support MC labels on the device but simply provide the ability to download them and rebuilt the MC on demand

f3sch · 2026-08-27T12:26:02Z

Do we need to duplicate the already existing ClusterLines math, can we not have only one of these classes?

f3sch · 2026-08-27T12:26:59Z

+  void prepareClusters(const TrackingParameters& trkParam, const int maxLayers) override
+  {
+    if (maxLayers < NLayers) { // only if former seeding vertexer is run
+      TimeFrame<NLayers>::prepareClusters(trkParam, maxLayers);
+    }
+  }
+  void allocateClusterSortStorage(const TrackingParameters& trkParam, const int maxLayers) override
+  {
+    if (maxLayers < NLayers) { // only if former seeding vertexer is run
+      TimeFrame<NLayers>::allocateClusterSortStorage(trkParam, maxLayers);
+    }
+  }


why defined in the header?

f3sch · 2026-08-27T12:28:37Z

+  std::bitset<NLayers> mUnsortedClustersUploaded{};
+  std::bitset<NLayers> mTrackingFrameInfoUploaded{};


These are not used anywhere?

f3sch · 2026-08-27T12:29:11Z

+  std::vector<o2::MCCompLabel> mLineLabelFlatHost;
+  int mNMembers{0};
+  Vertex* mDiamondDevice{nullptr};
+  bool mPersistentTablesLoaded{false};


Do we actually need this flag, can we not just add to the configuration the equivalent step?

f3sch · 2026-08-27T12:54:36Z

+  bool useParallelSeeding{false}; // use the GPU-oriented parallel seeding (computeVertices)
+  bool incrementalSeeding{false}; // CPU-only fast path for computeVertices


Why is this needed, we know by construction if we are running on CPU or GPU why do we need a configKeyVal?

f3sch · 2026-08-27T13:01:04Z

+  float lineMinPt = -1.f;
+  float fineZWindow = -1.f;
+  int fineMinDensity = 8;
+  float fineMaxDrift = -1.f;
+  float goodLineChi2Cut = 5.f;
+  float goodLinePtCut = 0.5f;
+  float goodContributorsSignificance = -1.f;
+  float duplicateZScale = -1.f;


In general all the config key values should be set to the Pb-Pb default values, we only really override them for the pp.

f3sch · 2026-08-27T13:03:59Z

+inline bool timeCompatible(const LineTime& a, const LineTime& b)
+{
+  return o2::gpu::GPUCommonMath::Abs(a.tc - b.tc) <= (a.th + b.th);
+}
+


if we anyways have LineTime class why not make this a function of that class?
Also If you were to use the already existing TimeEstBC class then we not even need that.

f3sch · 2026-08-27T13:04:55Z

+void computeDensities(const bounded_vector<float>& Z,
+                      const bounded_vector<LineTime>& T,
+                      const float zWindow,
+                      const bool incremental,
+                      bounded_vector<LineWindow>& win,
+                      bounded_vector<int>& density)
+{
+  const int m = static_cast<int>(Z.size());
+  int lo{0}, hi{0};
+  for (int k = 0; k < m; ++k) {
+    const float zk = Z[k];
+    if (incremental) {
+      while (lo < m && Z[lo] < zk - zWindow) {
+        ++lo;
+      }
+      while (hi < m && Z[hi] <= zk + zWindow) {
+        ++hi;
+      }
+    } else {
+      lo = static_cast<int>(std::lower_bound(Z.begin(), Z.end(), zk - zWindow) - Z.begin());
+      hi = static_cast<int>(std::upper_bound(Z.begin(), Z.end(), zk + zWindow) - Z.begin());
+    }
+    win[k] = LineWindow{lo, hi};
+    int count = 0;
+    for (int j = lo; j < hi; ++j) {
+      count += timeCompatible(T[k], T[j]);
+    }
+    density[k] = count;
+  }
+}
+
+void markLeftmostMaxima(const bounded_vector<int>& density,
+                        const bounded_vector<float>& Z,
+                        const bounded_vector<LineWindow>& win,
+                        const bool incremental,
+                        std::pmr::memory_resource* pool,
+                        bounded_vector<uint8_t>& isMax)
+{
+  const int m = static_cast<int>(density.size());
+  if (incremental) {
+    bounded_vector<int> maxDeque(m, pool);
+    int front{0}, back{0}, next{0};
+    for (int k = 0; k < m; ++k) {
+      while (next < win[k].hi) {
+        while (back > front && density[maxDeque[back - 1]] < density[next]) {
+          --back;
+        }
+        maxDeque[back++] = next++;
+      }
+      while (front < back && maxDeque[front] < win[k].lo) {
+        ++front;
+      }
+      isMax[k] = maxDeque[front] == k; // the window always contains k itself, so the deque is non-empty
+    }
+    return;
+  }
+  for (int k = 0; k < m; ++k) {
+    uint8_t best = 1;
+    for (int j = win[k].lo; j < win[k].hi; ++j) {
+      if (j == k) {
+        continue;
+      }
+      if (density[j] > density[k] || (density[j] == density[k] && Z[j] < Z[k])) {
+        best = 0;
+        break;
+      }
+    }
+    isMax[k] = best;
+  }
+}
+} // namespace


These functions are only called on the host, so incremental is always true?

f3sch · 2026-08-27T13:07:52Z

+        float deltaPhi{o2::gpu::CAMath::Abs(currentTracklet.phi - nextTracklet.phi)};
+        if (deltaPhi > o2::constants::math::PI) {
+          deltaPhi = o2::constants::math::TwoPI - deltaPhi;
+        }
+        if (deltaPhi > cellDeltaPhiCut) {
+          continue;
+        }


in MathUtils.h exist isPhiDifferenceBelow

cima22 · 2026-08-27T15:33:40Z

@shahor02 This algorithm should be tested, but I have some preliminary results.

On 50 simulated PbPb TFs at 50kHz:

On a PbPb TF with embedded pp collisions:

PbPb:

pp:

pp track efficiency:

shahor02 · 2026-08-27T15:43:41Z

Thanks, looks good! Note that there are conflicts in the PR.
So, will you be available to present it next Wed.?

cima22 · 2026-08-27T15:45:43Z

Thanks, looks good! Note that there are conflicts in the PR. So, will you be available to present it next Wed.?

Yes absolutely

fprino · 2026-08-27T17:36:59Z

Thanks, looks good! Note that there are conflicts in the PR. So, will you be available to present it next Wed.?

Yes absolutely

Sorry for replying late: next week we are considering to move the ITS WP2 meeting from Wednesday 2 september to Thursday 3 September, starting at 10 am.
Would this work for you (Gabriele, Ruben, Felix)?

f3sch · 2026-08-27T17:38:16Z

@fprino I will be on vacation from Thursday onwards.

f3sch · 2026-08-27T17:44:19Z

Some things from an initial test:
Timing:

[PROD][fschlepp@epn000 pvtest]$ python3 time.py gabriele/reco.log
Processing:
Original count: 925
Trimmed count:  649
Discarded:      138 low + 138 high
Total:          336405.15 ms (336.41 s)
Average:        518.34 ms
Std deviation:  194.22 ms

Vertex seeding:
Original count: 925
Trimmed count:  649
Discarded:      138 low + 138 high
Total:          60625.25 ms (60.63 s)
Average:        93.41 ms
Std deviation:  13.94 ms

Tracking:
Original count: 919
Trimmed count:  645
Discarded:      137 low + 137 high
Total:          231575.46 ms (231.58 s)
Average:        359.03 ms
Std deviation:  86.99 ms
[PROD][fschlepp@epn000 pvtest]$ python3 time.py dev/reco.log

Processing:
Original count: 925
Trimmed count:  649
Discarded:      138 low + 138 high
Total:          511572.46 ms (511.57 s)
Average:        788.25 ms
Std deviation:  239.93 ms

Vertex seeding:
Original count: 925
Trimmed count:  649
Discarded:      138 low + 138 high
Total:          222618.46 ms (222.62 s)
Average:        343.02 ms
Std deviation:  89.72 ms

Tracking:
Original count: 924
Trimmed count:  648
Discarded:      138 low + 138 high
Total:          240011.05 ms (240.01 s)
Average:        370.39 ms
Std deviation:  70.58 ms

note this timing from MI50 EPNs. the new vertexing seems to be 3.7 times faster which in total means a speedup of 1.5 in total processing time (-34.2%).

23 kHz Pb-Pb.
red:dev, blue:this PR, green:this PR but relaxing ITSVertexerParam.clusterContributorsCut from 3 to 2:
pv_cont.pdf
There seems be recovery of PVs below 80 contributors and a slight degradation for higher mult vertices.

black:dev, red:this PR, blue:this PR but relaxing ITSVertexerParam.clusterContributorsCut from 3 to 2:
K0sFitOverlay.pdf
K0sFitVsPt.pdf
K0s yield decreases by 0.2% although this is dominated by the 1-1.5 GeV/c bin.

Note that there is currently no second iteration for UPC flagging.

shahor02 · 2026-08-27T18:05:39Z

Then let's have it presented at the wp12/13 meeting next Wed. 15:30.
@cima22 Was suppressLowMultDebris=13 added to pp settings intentionally? Usually, we don't remove debris there.
@f3sch thanks, looks good. I assume for the UPC you left default settings?

f3sch · 2026-08-27T18:09:29Z

@f3sch thanks, looks good. I assume for the UPC you left default settings?

@shahor02 there is only one iteration done in this PR, if you mean the one from dev then yes.

shahor02 · 2026-08-27T18:29:55Z

OK, I saw https://github.com/AliceO2Group/AliceO2/pull/15733/changes#diff-27fc8aaeb7769a496f838182d30a759ba63b3163bc7f458dd9bbbc895d3c7211R258-R260, but then the ITSVertexerParam are irrelevant in that mode, and there is no ITS trackfinding UPC iteration?
If the answer is long, don't waste time, will learn at the presentation.

f3sch · 2026-08-27T19:09:24Z

OK, I saw https://github.com/AliceO2Group/AliceO2/pull/15733/changes#diff-27fc8aaeb7769a496f838182d30a759ba63b3163bc7f458dd9bbbc895d3c7211R258-R260, but then the ITSVertexerParam are irrelevant in that mode, and there is no ITS trackfinding UPC iteration?

There is no track finding since no ROFs are flagged by a vertex carrying the UPC flag in the UPC mask, all ROFs are automatically masked in UPC iteration

cima22 · 2026-08-28T07:13:15Z

Then let's have it presented at the wp12/13 meeting next Wed. 15:30.

@cima22 Was suppressLowMultDebris=13 added to pp settings intentionally? Usually, we don't remove debris there.

@f3sch thanks, looks good. I assume for the UPC you left default settings?

Hi, sorry today I am travelling so I won't be very responsive. I didn't put yet the UPC iteration in the PR, I will test one in the next days.

suppressLowMultDebris=13 slipped when copying the parameters, but in any case it is not used by this algorithm by default, thanks for noticing.

cima22 requested review from f3sch, fprino, mconcas, mcoquet642 and shahor02 as code owners August 27, 2026 11:12

alibuild mentioned this pull request Aug 27, 2026

Please consider the following formatting changes to #15733 cima22/AliceO2#10

Closed

ITS: new CPU + GPU seeding vertexer

06b0b02

Adds a seeding vertexer that runs as a prepended tracker pass (diamond trackleting -> cells -> lines -> parallel seeding), on both the CPU and GPU traits, replacing the per-ROF CPU vertexer for the seeding step.

cima22 force-pushed the parallel-vertexer branch from 5c4ab24 to 06b0b02 Compare August 27, 2026 11:18

f3sch reviewed Aug 27, 2026

View reviewed changes

		std::bitset<NLayers> mUnsortedClustersUploaded{};
		std::bitset<NLayers> mTrackingFrameInfoUploaded{};

		bool useParallelSeeding{false}; // use the GPU-oriented parallel seeding (computeVertices)
		bool incrementalSeeding{false}; // CPU-only fast path for computeVertices

Uh oh!

Conversation

cima22 commented Aug 27, 2026

Uh oh!

shahor02 commented Aug 27, 2026

Uh oh!

f3sch left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

cima22 commented Aug 27, 2026

On 50 simulated PbPb TFs at 50kHz:

On a PbPb TF with embedded pp collisions:

PbPb:

pp:

pp track efficiency:

Uh oh!

shahor02 commented Aug 27, 2026

Uh oh!

cima22 commented Aug 27, 2026

Uh oh!

fprino commented Aug 27, 2026

Uh oh!

f3sch commented Aug 27, 2026

Uh oh!

f3sch commented Aug 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

shahor02 commented Aug 27, 2026

Uh oh!

f3sch commented Aug 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

shahor02 commented Aug 27, 2026

Uh oh!

f3sch commented Aug 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

cima22 commented Aug 28, 2026

Uh oh!

Reviewers

Assignees

Labels

Milestone

Development

Uh oh!

4 participants

f3sch commented Aug 27, 2026 •

edited

Loading

f3sch commented Aug 27, 2026 •

edited

Loading

f3sch commented Aug 27, 2026 •

edited

Loading